#include "kernel_operator.h"
using namespace AscendC;
constexpr int32_t BUFFER_NUM = 2;

class KernelLerp
{
public:
  __aicore__ inline KernelLerp() {}
  __aicore__ inline void Init(GM_ADDR start, GM_ADDR end, GM_ADDR weight, GM_ADDR y,
                              uint32_t datatypeclass, uint32_t broadcast,
                              uint32_t blockLength, uint32_t startlastDim,
                              uint32_t tileNum, uint32_t tileLength,
                              uint32_t lasttileLength, uint32_t formerNum,
                              uint32_t formerLength, uint32_t formertileNum,
                              uint32_t formertileLength,
                              uint32_t formerlasttileLength, uint32_t tailNum,
                              uint32_t tailLength, uint32_t tailtileNum,
                              uint32_t tailtileLength, uint32_t taillasttileLength,
                              uint32_t tilingKey, uint32_t startLength, uint32_t endLength,
                              uint32_t weightLength)
  {
    ASSERT(GetBlockNum() != 0 && "block dim can not be zero!");

    uint32_t gen_class = 0;
    if (weightLength == 1)
    {
      gen_class = 1; // weight 常数
    }
    else if (startLength / weightLength > 1)
    {
      gen_class = 2;
    } // weight tensor且需要广播

    this->broadcast = broadcast;
    this->gen_class = gen_class;
    this->endLength = endLength;
    this->datatypeclass = datatypeclass;
    this->tilingKey = tilingKey;
    this->weightLength = weightLength;
    this->startlastDim = startlastDim;

    if (tilingKey == 1)
    { // 对齐场景
      this->blockLength = blockLength;
      uint32_t offset = this->blockLength * GetBlockIdx();
      this->tileNum = tileNum;
      this->lasttileLength = lasttileLength;
      this->tileLength = tileLength / BUFFER_NUM;

      if (this->gen_class == 0)
      {

        if (this->broadcast)
        {
          endGm.SetGlobalBuffer((__gm__ DTYPE_END *)end, this->endLength);
        }
        else
        {
          endGm.SetGlobalBuffer((__gm__ DTYPE_END *)end + offset,
                                this->blockLength);
        }
        startGm.SetGlobalBuffer((__gm__ DTYPE_START *)start + offset,
                                this->blockLength);
        weightGm.SetGlobalBuffer((__gm__ DTYPE_WEIGHT *)weight + offset,
                                 this->blockLength);
        yGm.SetGlobalBuffer((__gm__ DTYPE_Y *)y + offset,
                            this->blockLength);
      }
      else
      {
        if (this->broadcast)
        {
          endGm.SetGlobalBuffer((__gm__ DTYPE_END *)end, this->endLength);
        }
        else
        {
          endGm.SetGlobalBuffer((__gm__ DTYPE_END *)end + offset,
                                this->blockLength);
        }
        startGm.SetGlobalBuffer((__gm__ DTYPE_START *)start + offset,
                                this->blockLength);
        weightGm.SetGlobalBuffer((__gm__ DTYPE_WEIGHT *)weight, this->weightLength);
        yGm.SetGlobalBuffer((__gm__ DTYPE_Y *)y + offset,
                            this->blockLength);
      }
    }

    if (tilingKey == 2)
    { // 不均分
      this->formerNum = formerNum;
      this->formerLength = formerLength;
      this->formertileNum = formertileNum;
      this->formertileLength = formertileLength;
      this->formerlasttileLength = formerlasttileLength;

      this->tailNum = tailNum;
      this->tailLength = tailLength;
      this->tailtileNum = tailtileNum;
      this->tailtileLength = tailtileLength;
      this->taillasttileLength = taillasttileLength;

      if (GetBlockIdx() < this->formerNum)
      { // 大块
        this->tileLength = this->formertileLength / BUFFER_NUM;
        this->lasttileLength = this->formerlasttileLength;
        this->tileNum = this->formertileNum * BUFFER_NUM;
        uint32_t offset = this->formerLength * GetBlockIdx();

        if (this->gen_class == 0)
        {

          if (this->broadcast)
          {
            endGm.SetGlobalBuffer((__gm__ DTYPE_END *)end, this->endLength);
          }
          else
          {
            endGm.SetGlobalBuffer((__gm__ DTYPE_END *)end + offset,
                                  this->formerLength);
          }
          startGm.SetGlobalBuffer((__gm__ DTYPE_START *)start + offset,
                                  this->formerLength);
          weightGm.SetGlobalBuffer((__gm__ DTYPE_WEIGHT *)weight + offset,
                                   this->formerLength);
          yGm.SetGlobalBuffer((__gm__ DTYPE_WEIGHT *)y + offset,
                              this->formerLength);
        }
        else
        {
          if (this->broadcast)
          {
            endGm.SetGlobalBuffer((__gm__ DTYPE_END *)end, this->endLength);
          }
          else
          {
            endGm.SetGlobalBuffer((__gm__ DTYPE_END *)end + offset,
                                  this->formerLength);
          }
          startGm.SetGlobalBuffer((__gm__ DTYPE_START *)start + offset,
                                  this->formerLength);
          weightGm.SetGlobalBuffer((__gm__ DTYPE_WEIGHT *)weight, this->weightLength);
          this->weightScalar = weightGm.GetValue(0);
          yGm.SetGlobalBuffer((__gm__ DTYPE_WEIGHT *)y + offset,
                              this->formerLength);
        }
      }
      else
      { // 小块
        this->tileLength = this->tailtileLength / BUFFER_NUM;
        this->lasttileLength = this->taillasttileLength;
        this->tileNum = this->tailtileNum * BUFFER_NUM;
        uint32_t offset = this->formerLength * this->formerNum + this->tailLength * (GetBlockIdx() - this->formerNum);
        if (this->gen_class == 0)
        {
          if (this->broadcast)
          {
            endGm.SetGlobalBuffer((__gm__ DTYPE_END *)end, this->endLength);
          }
          else
          {
            endGm.SetGlobalBuffer((__gm__ DTYPE_END *)end + offset,
                                  this->tailLength);
          }
          startGm.SetGlobalBuffer((__gm__ DTYPE_START *)start + offset,
                                  this->tailLength);
          weightGm.SetGlobalBuffer((__gm__ DTYPE_WEIGHT *)weight + offset,
                                   this->tailLength);
          yGm.SetGlobalBuffer((__gm__ DTYPE_WEIGHT *)y + offset,
                              this->tailLength);
        }
        else
        {
          if (this->broadcast)
          {
            endGm.SetGlobalBuffer((__gm__ DTYPE_END *)end, this->endLength);
          }
          else
          {
            endGm.SetGlobalBuffer((__gm__ DTYPE_END *)end + offset,
                                  this->tailLength);
          }
          startGm.SetGlobalBuffer((__gm__ DTYPE_START *)start + offset,
                                  this->tailLength);
          weightGm.SetGlobalBuffer((__gm__ DTYPE_WEIGHT *)weight, weightLength);
          this->weightScalar = weightGm.GetValue(0);
          yGm.SetGlobalBuffer((__gm__ DTYPE_WEIGHT *)y + offset,
                              this->tailLength);
        }
      }
    } // 不均分

    pipe.InitBuffer(inQueueStart, BUFFER_NUM, this->tileLength * sizeof(DTYPE_START));
    pipe.InitBuffer(inQueueEnd, BUFFER_NUM, this->tileLength * sizeof(DTYPE_END));
    pipe.InitBuffer(inQueueWeight, BUFFER_NUM, this->tileLength * sizeof(DTYPE_WEIGHT));
    pipe.InitBuffer(outQueueY, BUFFER_NUM, this->tileLength * sizeof(DTYPE_Y));
    pipe.InitBuffer(tmp1, this->tileLength * sizeof(float));
    pipe.InitBuffer(tmp2, this->tileLength * sizeof(float));

  } // Init

  __aicore__ inline void Process()
  {
    int32_t loopCount = this->tileNum * BUFFER_NUM;
    for (int32_t i = 0; i < loopCount; i++)
    {
      CopyIn(i);
      Compute(i);
      CopyOut(i);
    }
  } // Process()

private:
  __aicore__ inline void CopyIn(int32_t progress)
  {
    if (this->gen_class == 0)
    {
      LocalTensor<DTYPE_START> startLocal = inQueueStart.AllocTensor<DTYPE_START>();
      LocalTensor<DTYPE_END> endLocal = inQueueEnd.AllocTensor<DTYPE_END>();
      LocalTensor<DTYPE_WEIGHT> weightLocal = inQueueWeight.AllocTensor<DTYPE_WEIGHT>();
      if (this->broadcast == 1)
      {
        int32_t start_e;
        if ((progress == (this->tileNum * BUFFER_NUM - 2)) ||
            (progress == (this->tileNum * BUFFER_NUM - 1)))
        {
          DataCopy(startLocal,
                   startGm[(progress - 2) * (this->tileLength) + this->lasttileLength],
                   (this->tileLength));
          DataCopy(weightLocal,
                   weightGm[((progress - 2) * (this->tileLength) + this->lasttileLength)],
                   (this->tileLength));
          if (this->tilingKey == 1)
          {
            start_e = GetBlockIdx() * this->blockLength + (progress - 2) * (this->tileLength) + this->lasttileLength;
          }
          else
          {
            start_e = ((GetBlockIdx() < this->formerNum) ? (this->formerLength * GetBlockIdx()) : (this->formerLength * this->formerNum + (GetBlockIdx() - this->formerNum) * this->tailLength)) + (progress - 2) * (this->tileLength) + this->lasttileLength;
          }
        }
        else
        {
          DataCopy(startLocal, startGm[progress * (this->tileLength)], (this->tileLength));
          DataCopy(weightLocal, weightGm[(progress * this->tileLength)], (this->tileLength));
          if (this->tilingKey == 1)
          {
            start_e = GetBlockIdx() * this->blockLength + progress * this->tileLength;
          }
          else
          {
            start_e = ((GetBlockIdx() < this->formerNum) ? (this->formerLength * GetBlockIdx()) : (this->formerLength * this->formerNum + (GetBlockIdx() - this->formerNum) * this->tailLength)) + progress * this->tileLength;
          }
        }
        int32_t gmStart_e = start_e % this->endLength;
        int32_t copyLength = this->tileLength;
        for (int32_t i = 0; i < copyLength; i++)
        {
          endLocal.SetValue(i, endGm.GetValue((gmStart_e + i) % this->endLength));
        }
      }
      else if (this->broadcast == 2)
      {
        int32_t start_e;
        if ((progress == (this->tileNum * BUFFER_NUM - 2)) ||
            (progress == (this->tileNum * BUFFER_NUM - 1)))
        {
          DataCopy(startLocal,
                   startGm[(progress - 2) * (this->tileLength) + this->lasttileLength],
                   (this->tileLength));
          DataCopy(weightLocal,
                   weightGm[((progress - 2) * (this->tileLength) + this->lasttileLength)],
                   (this->tileLength));
          if (this->tilingKey == 1)
          {
            start_e = GetBlockIdx() * this->blockLength + (progress - 2) * (this->tileLength) + this->lasttileLength;
          }
          else
          {
            start_e = ((GetBlockIdx() < this->formerNum) ? (this->formerLength * GetBlockIdx()) : (this->formerLength * this->formerNum + (GetBlockIdx() - this->formerNum) * this->tailLength)) + (progress - 2) * (this->tileLength) + this->lasttileLength;
          }
        }
        else
        {
          DataCopy(startLocal, startGm[progress * (this->tileLength)], (this->tileLength));
          DataCopy(weightLocal, weightGm[(progress * this->tileLength)], (this->tileLength));
          if (this->tilingKey == 1)
          {
            start_e = GetBlockIdx() * this->blockLength + progress * this->tileLength;
          }
          else
          {
            start_e = ((GetBlockIdx() < this->formerNum) ? (this->formerLength * GetBlockIdx()) : (this->formerLength * this->formerNum + (GetBlockIdx() - this->formerNum) * this->tailLength)) + progress * this->tileLength;
          }
        }
        int32_t gmStart_e = start_e;
        int32_t copyLength = this->tileLength;
        for (int32_t i = 0; i < copyLength; i++)
        {
          endLocal.SetValue(i, endGm.GetValue((gmStart_e + i) / this->startlastDim));
        }
      }
      else
      {
        if ((progress == (this->tileNum * BUFFER_NUM - 2)) ||
            (progress == (this->tileNum * BUFFER_NUM - 1)))
        {
          DataCopy(startLocal,
                   startGm[(progress - 2) * (this->tileLength) + this->lasttileLength],
                   (this->tileLength));
          DataCopy(endLocal,
                   endGm[(progress - 2) * (this->tileLength) + this->lasttileLength],
                   (this->tileLength));
          DataCopy(weightLocal,
                   weightGm[((progress - 2) * (this->tileLength) + this->lasttileLength)],
                   (this->tileLength));
        }
        else
        {
          DataCopy(startLocal, startGm[progress * (this->tileLength)], (this->tileLength));
          DataCopy(endLocal, endGm[progress * (this->tileLength)], (this->tileLength));
          DataCopy(weightLocal, weightGm[(progress * this->tileLength)], (this->tileLength));
        }
      }
      inQueueStart.EnQue(startLocal);
      inQueueEnd.EnQue(endLocal);
      inQueueWeight.EnQue(weightLocal);
    }
    else if (gen_class == 1)
    {
      LocalTensor<DTYPE_START> startLocal = inQueueStart.AllocTensor<DTYPE_START>();
      LocalTensor<DTYPE_END> endLocal = inQueueEnd.AllocTensor<DTYPE_END>();
      LocalTensor<DTYPE_WEIGHT> weightLocal = inQueueWeight.AllocTensor<DTYPE_WEIGHT>();
      if (this->broadcast == 1)
      {
        int32_t start_e;
        if ((progress == (this->tileNum * BUFFER_NUM - 2)) ||
            (progress == (this->tileNum * BUFFER_NUM - 1)))
        {
          DataCopy(startLocal,
                   startGm[(progress - 2) * (this->tileLength) + this->lasttileLength],
                   (this->tileLength));
          DataCopy(weightLocal, weightGm[0], (this->tileLength));
          if (this->tilingKey == 1)
          {
            start_e = GetBlockIdx() * this->blockLength + (progress - 2) * (this->tileLength) + this->lasttileLength;
          }
          else
          {
            start_e = ((GetBlockIdx() < this->formerNum) ? (this->formerLength * GetBlockIdx()) : (this->formerLength * this->formerNum + (GetBlockIdx() - this->formerNum) * this->tailLength)) + (progress - 2) * (this->tileLength) + this->lasttileLength;
          }
        }
        else
        {
          DataCopy(startLocal, startGm[progress * (this->tileLength)], (this->tileLength));
          DataCopy(weightLocal, weightGm[0], (this->tileLength));
          if (this->tilingKey == 1)
          {
            start_e = GetBlockIdx() * this->blockLength + progress * this->tileLength;
          }
          else
          {
            start_e = ((GetBlockIdx() < this->formerNum) ? (this->formerLength * GetBlockIdx()) : (this->formerLength * this->formerNum + (GetBlockIdx() - this->formerNum) * this->tailLength)) + progress * this->tileLength;
          }
        }
        int32_t gmStart_e = start_e % this->endLength;
        int32_t copyLength = this->tileLength;
        for (int32_t i = 0; i < copyLength; i++)
        {
          endLocal.SetValue(i, endGm.GetValue((gmStart_e + i) % this->endLength));
        }
      }
      else if (this->broadcast == 2)
      {
        int32_t start_e;
        if ((progress == (this->tileNum * BUFFER_NUM - 2)) ||
            (progress == (this->tileNum * BUFFER_NUM - 1)))
        {
          DataCopy(startLocal,
                   startGm[(progress - 2) * (this->tileLength) + this->lasttileLength],
                   (this->tileLength));
          DataCopy(weightLocal, weightGm[0], (this->tileLength));
          if (this->tilingKey == 1)
          {
            start_e = GetBlockIdx() * this->blockLength + (progress - 2) * (this->tileLength) + this->lasttileLength;
          }
          else
          {
            start_e = ((GetBlockIdx() < this->formerNum) ? (this->formerLength * GetBlockIdx()) : (this->formerLength * this->formerNum + (GetBlockIdx() - this->formerNum) * this->tailLength)) + (progress - 2) * (this->tileLength) + this->lasttileLength;
          }
        }
        else
        {
          DataCopy(startLocal, startGm[progress * (this->tileLength)], (this->tileLength));
          DataCopy(weightLocal, weightGm[0], (this->tileLength));
          if (this->tilingKey == 1)
          {
            start_e = GetBlockIdx() * this->blockLength + progress * this->tileLength;
          }
          else
          {
            start_e = ((GetBlockIdx() < this->formerNum) ? (this->formerLength * GetBlockIdx()) : (this->formerLength * this->formerNum + (GetBlockIdx() - this->formerNum) * this->tailLength)) + progress * this->tileLength;
          }
        }
        int32_t gmStart_e = start_e;
        int32_t copyLength = this->tileLength;
        for (int32_t i = 0; i < copyLength; i++)
        {
          endLocal.SetValue(i, endGm.GetValue((gmStart_e + i) / this->startlastDim));
        }
      }
      else
      {
        if ((progress == (this->tileNum * BUFFER_NUM - 2)) ||
            (progress == (this->tileNum * BUFFER_NUM - 1)))
        {
          DataCopy(startLocal,
                   startGm[(progress - 2) * (this->tileLength) + this->lasttileLength],
                   (this->tileLength));
          DataCopy(endLocal,
                   endGm[(progress - 2) * (this->tileLength) + this->lasttileLength],
                   (this->tileLength));
          DataCopy(weightLocal, weightGm[0], (this->tileLength));
        }
        else
        {
          DataCopy(startLocal, startGm[progress * (this->tileLength)], (this->tileLength));
          DataCopy(endLocal, endGm[progress * (this->tileLength)], (this->tileLength));
          DataCopy(weightLocal, weightGm[0], (this->tileLength));
        }
      }
      inQueueStart.EnQue(startLocal);
      inQueueEnd.EnQue(endLocal);
      inQueueWeight.EnQue(weightLocal);
    }
    else if (gen_class == 2)
    {
      LocalTensor<DTYPE_START> startLocal = inQueueStart.AllocTensor<DTYPE_START>();
      LocalTensor<DTYPE_END> endLocal = inQueueEnd.AllocTensor<DTYPE_END>();
      LocalTensor<DTYPE_WEIGHT> weightLocal = inQueueWeight.AllocTensor<DTYPE_WEIGHT>();
      int32_t start_w;
      if ((progress == (this->tileNum * BUFFER_NUM - 2)) ||
          (progress == (this->tileNum * BUFFER_NUM - 1)))
      {
        DataCopy(startLocal,
                 startGm[(progress - 2) * (this->tileLength) + this->lasttileLength],
                 (this->tileLength));
        if (this->tilingKey == 1)
        {
          start_w = GetBlockIdx() * this->blockLength + (progress - 2) * (this->tileLength) + this->lasttileLength;
        }
        else
        {
          start_w = ((GetBlockIdx() < this->formerNum) ? (this->formerLength * GetBlockIdx()) : (this->formerLength * this->formerNum + (GetBlockIdx() - this->formerNum) * this->tailLength)) + (progress - 2) * (this->tileLength) + this->lasttileLength;
        }
      }
      else
      {
        DataCopy(startLocal, startGm[progress * (this->tileLength)], (this->tileLength));
        if (this->tilingKey == 1)
        {
          start_w = GetBlockIdx() * this->blockLength + progress * this->tileLength;
        }
        else
        {
          start_w = ((GetBlockIdx() < this->formerNum) ? (this->formerLength * GetBlockIdx()) : (this->formerLength * this->formerNum + (GetBlockIdx() - this->formerNum) * this->tailLength)) + progress * this->tileLength;
        }
      }
      int32_t gmStart_w = start_w;
      int32_t copyLength = this->tileLength;
      for (int32_t i = 0; i < copyLength; i++)
      {
        weightLocal.SetValue(i, weightGm.GetValue((gmStart_w + i) /this->weightLength));
      }

      if (this->broadcast == 1)
      {
        int32_t start_e;
        if ((progress == (this->tileNum * BUFFER_NUM - 2)) ||
            (progress == (this->tileNum * BUFFER_NUM - 1)))
        {
          if (this->tilingKey == 1)
          {
            start_e = GetBlockIdx() * this->blockLength + (progress - 2) * (this->tileLength) + this->lasttileLength;
          }
          else
          {
            start_e = ((GetBlockIdx() < this->formerNum) ? (this->formerLength * GetBlockIdx()) : (this->formerLength * this->formerNum + (GetBlockIdx() - this->formerNum) * this->tailLength)) + (progress - 2) * (this->tileLength) + this->lasttileLength;
          }
        }
        else
        {
          if (this->tilingKey == 1)
          {
            start_e = GetBlockIdx() * this->blockLength + progress * this->tileLength;
          }
          else
          {
            start_e = ((GetBlockIdx() < this->formerNum) ? (this->formerLength * GetBlockIdx()) : (this->formerLength * this->formerNum + (GetBlockIdx() - this->formerNum) * this->tailLength)) + progress * this->tileLength;
          }
        }
        int32_t gmStart_e = start_e % this->endLength;
        int32_t copyLength = this->tileLength;
        for (int32_t i = 0; i < copyLength; i++)
        {
          endLocal.SetValue(i, endGm.GetValue((gmStart_e + i) % this->endLength));
        }
      }
      else if (this->broadcast == 2)
      {
        int32_t start_e;
        if ((progress == (this->tileNum * BUFFER_NUM - 2)) ||
            (progress == (this->tileNum * BUFFER_NUM - 1)))
        {
          if (this->tilingKey == 1)
          {
            start_e = GetBlockIdx() * this->blockLength + (progress - 2) * (this->tileLength) + this->lasttileLength;
          }
          else
          {
            start_e = ((GetBlockIdx() < this->formerNum) ? (this->formerLength * GetBlockIdx()) : (this->formerLength * this->formerNum + (GetBlockIdx() - this->formerNum) * this->tailLength)) + (progress - 2) * (this->tileLength) + this->lasttileLength;
          }
        }
        else
        {
          if (this->tilingKey == 1)
          {
            start_e = GetBlockIdx() * this->blockLength + progress * this->tileLength;
          }
          else
          {
            start_e = ((GetBlockIdx() < this->formerNum) ? (this->formerLength * GetBlockIdx()) : (this->formerLength * this->formerNum + (GetBlockIdx() - this->formerNum) * this->tailLength)) + progress * this->tileLength;
          }
        }
        int32_t gmStart_e = start_e;
        int32_t copyLength = this->tileLength;
        for (int32_t i = 0; i < copyLength; i++)
        {
          endLocal.SetValue(i, endGm.GetValue((gmStart_e + i) / this->startlastDim));
        }
      }
      else
      {
        if ((progress == (this->tileNum * BUFFER_NUM - 2)) ||
            (progress == (this->tileNum * BUFFER_NUM - 1)))
        {
          DataCopy(endLocal,
                   endGm[(progress - 2) * (this->tileLength) + this->lasttileLength],
                   (this->tileLength));
        }
        else
        {
          DataCopy(endLocal, endGm[progress * (this->tileLength)], (this->tileLength));
        }
      }
      inQueueStart.EnQue(startLocal);
      inQueueEnd.EnQue(endLocal);
      inQueueWeight.EnQue(weightLocal);
    } // class
  } // CopyIn()

  __aicore__ inline void Compute(int32_t progress)
  {
    LocalTensor<DTYPE_START> startLocal = inQueueStart.DeQue<DTYPE_START>();
    LocalTensor<DTYPE_END> endLocal = inQueueEnd.DeQue<DTYPE_END>();
    LocalTensor<DTYPE_WEIGHT> weightLocal = inQueueWeight.DeQue<DTYPE_WEIGHT>();
    LocalTensor<DTYPE_Y> yLocal = outQueueY.AllocTensor<DTYPE_Y>();
    auto tmp1Local = tmp1.Get<float>();
    auto tmp2Local = tmp2.Get<float>();

    if (this->gen_class == 0 || this->gen_class == 2)
    {
      if (this->datatypeclass == 0)
      {
        Cast(tmp1Local, startLocal, RoundMode::CAST_NONE, this->tileLength);
        Cast(tmp2Local, endLocal, RoundMode::CAST_NONE, this->tileLength);
        Sub(tmp2Local, tmp2Local, tmp1Local, this->tileLength);
        Cast(tmp1Local, weightLocal, RoundMode::CAST_NONE, this->tileLength);
        Mul(tmp2Local, tmp2Local, tmp1Local, this->tileLength);
        Cast(tmp1Local, startLocal, RoundMode::CAST_NONE, this->tileLength);
        Add(tmp2Local, tmp2Local, tmp1Local, this->tileLength);
        Cast(yLocal, tmp2Local, RoundMode::CAST_NONE, this->tileLength);
      }
      else
      {
        Sub(endLocal, endLocal, startLocal, this->tileLength);
        Mul(yLocal, endLocal, weightLocal, this->tileLength);
        Add(yLocal, yLocal, startLocal, this->tileLength);
      }
    }
    else if (this->gen_class == 1)
    {
      if (this->datatypeclass == 0)
      {
        Cast(tmp1Local, startLocal, RoundMode::CAST_NONE, this->tileLength);
        Cast(tmp2Local, endLocal, RoundMode::CAST_NONE, this->tileLength);
        Sub(tmp2Local, tmp2Local, tmp1Local, this->tileLength);
        float weight = this->weightScalar;
        Muls(tmp2Local, tmp2Local, weight, this->tileLength);
        Add(tmp2Local, tmp2Local, tmp1Local, this->tileLength);
        Cast(yLocal, tmp2Local, RoundMode::CAST_NONE, this->tileLength);
      }
      else
      {
        Sub(endLocal, endLocal, startLocal, this->tileLength);
        Muls(yLocal, endLocal, this->weightScalar, this->tileLength);
        Add(yLocal, yLocal, startLocal, this->tileLength);
      }
    }
    outQueueY.EnQue<DTYPE_Y>(yLocal);
    inQueueStart.FreeTensor(startLocal);
    inQueueEnd.FreeTensor(endLocal);
    inQueueWeight.FreeTensor(weightLocal);
  } // Compute

  __aicore__ inline void CopyOut(int32_t progress)
  {
    LocalTensor<DTYPE_Y> yLocal = outQueueY.DeQue<DTYPE_Y>();
    if ((progress == (this->tileNum * BUFFER_NUM - 2)) ||
        (progress == (this->tileNum * BUFFER_NUM - 1)))
    {
      DataCopy(
          yGm[(progress - 2) * (this->tileLength) + this->lasttileLength],
          yLocal, (this->tileLength));
    }

    else
    {
      DataCopy(yGm[progress * (this->tileLength)], yLocal,
               (this->tileLength));
    }
    outQueueY.FreeTensor(yLocal);
  } // CopyOut()

private:
  TPipe pipe;
  TQue<QuePosition::VECIN, BUFFER_NUM> inQueueStart, inQueueEnd, inQueueWeight;
  TQue<QuePosition::VECOUT, BUFFER_NUM> outQueueY;
  TBuf<QuePosition::VECCALC> tmp1, tmp2;
  GlobalTensor<DTYPE_START> startGm;
  GlobalTensor<DTYPE_END> endGm;
  GlobalTensor<DTYPE_WEIGHT> weightGm;
  GlobalTensor<DTYPE_Y> yGm;
  uint32_t blockLength;
  uint32_t tileNum;
  uint32_t tileLength;
  uint32_t lasttileLength;
  uint32_t formerNum;
  uint32_t formerLength;
  uint32_t formertileNum;
  uint32_t formertileLength;
  uint32_t datatypeclass;
  uint32_t formerlasttileLength;
  uint32_t tailNum;
  uint32_t gen_class;
  uint32_t tailLength;
  uint32_t tailtileNum;
  uint32_t tailtileLength;
  uint32_t taillasttileLength;
  uint32_t broadcast;
  uint32_t endLength;
  uint32_t weightLength;
  uint32_t tilingKey;
  uint32_t startlastDim;
  DTYPE_WEIGHT weightScalar;
};

extern "C" __global__ __aicore__ void lerp(GM_ADDR start, GM_ADDR end, GM_ADDR weight, GM_ADDR y, GM_ADDR workspace, GM_ADDR tiling)
{
  GET_TILING_DATA(tilingData, tiling);
  // TODO: user kernel impl
  KernelLerp op;
  uint32_t tilingKey = 1;
  if (TILING_KEY_IS(1))
  {
    tilingKey = 1;
  }
  else if (TILING_KEY_IS(2))
  {
    tilingKey = 2;
  }
  else
  {
    tilingKey = 1;
  }
  op.Init(start, end, weight, y, tilingData.datatypeclass, tilingData.broadcast,
          tilingData.blockLength, tilingData.startlastDim,
          tilingData.tileNum, tilingData.tileLength,
          tilingData.lasttileLength, tilingData.formerNum,
          tilingData.formerLength, tilingData.formertileNum,
          tilingData.formertileLength, tilingData.formerlasttileLength,
          tilingData.tailNum, tilingData.tailLength, tilingData.tailtileNum,
          tilingData.tailtileLength, tilingData.taillasttileLength,
          tilingKey, tilingData.startLength, tilingData.endLength, tilingData.weightLength);
  op.Process();
}

